// Copyright (C) 2019. Huawei Technologies Co., Ltd. All rights reserved.

// Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"),
// to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense,
// and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions:

// The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software.

// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE
// WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR
// COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.





__kernel void transpose_nchw_0132(const int ih_str, const int iw_str, const int ih_off, const int iw_off, const int oh_str, const int ow_str, 
    const int oh_off, const int ow_off, const int oh, const int bx, const int by, __global const T* input, __global T* output) {

    int idx = get_global_id(0);
    int idy = get_global_id(1);
    int idz = get_global_id(2);
    if(idx >= bx || idy >= by) return;
    T4 val;
    int idx4 = idx << 2;
    int in_off = (idz * ih_str + idy + ih_off) * iw_str + idx4 + iw_off;
    val = vload4(0, input + in_off);
    int out_off = (idz * oh_str + idx4 + oh_off) * ow_str + idy + ow_off;
    output[out_off] = val.x;
    if(idx4 + 1 < oh) output[out_off + ow_str]     = val.y;
    if(idx4 + 2 < oh) output[out_off + ow_str * 2] = val.z;
    if(idx4 + 3 < oh) output[out_off + ow_str * 3] = val.w;
}
